You write custom CUDA kernels to replace pytorch operators in given architecture to get speedups. You have complete freedom to choose set of operators you want to replace. You may make the decision to replace some operators with custom CUDA kernels and leave others unchanged. You may replace multiple operators with custom implementations, consider operator fusion opportunities (combining multiple operators into a single kernel, for example, combining matmul+relu), or algorithmic changes (such as online softmax). You are only limited by your imagination.

**SPECIAL INSTRUCTIONS FOR MANHATTAN + MSE FUSION:**

When implementing Manhattan Distance + MSE Loss fusion, you MUST implement the following optimized strategy:

1. **FUSION ARCHITECTURE**: Combine Manhattan distance and MSE loss computation in a single kernel:
   - Compute both Manhattan distance (L1) and MSE loss (L2) simultaneously
   - Process both metrics in the same memory access pattern
   - Eliminate intermediate tensor storage for maximum efficiency
   - Store both Manhattan distances and MSE losses as separate outputs

2. **DUAL METRIC COMPUTATION**: Implement both distance metrics in parallel:
   - Manhattan Distance: Σ|x_i - y_i|
   - MSE Loss: Σ(x_i - y_i)² / feature_dim
   - Compute both using the same difference values
   - Use shared memory to store partial results for both metrics

3. **FLOAT4 VECTORIZATION**: Use float4 vectorization for maximum memory bandwidth utilization:
   - Process 4 elements simultaneously using float4 loads
   - Compute both Manhattan and MSE for all 4 components in parallel
   - Handle remaining elements with scalar processing
   - Minimize memory access through vectorized operations

4. **WARP-LEVEL OPTIMIZATION**: Use warp-level processing for maximum performance:
   - Each block processes one sample from the batch
   - Use 8 warps per block (256 threads) for optimal GPU utilization
   - Use __shfl_down_sync for efficient warp-level reduction of both metrics
   - Divide feature dimensions among warps for parallel processing

5. **DUAL REDUCTION STRATEGY**: Implement two distinct reduction approaches:
   - **Fast Mode**: Use shared memory + atomic operations for both metrics
   - **Efficient Mode**: Use two-level reduction (partial sums + CPU final reduction)
   - Both modes must handle Manhattan and MSE simultaneously
   - Use mode parameter to select between reduction strategies

6. **SHARED MEMORY PATTERN**: Use efficient shared memory organization:
cpp
// For dual metric reduction across warps
extern __shared__ float shared_data[];
if (lane_id == 0) {
    shared_data[warp_id] = warp_manhattan;        // First half: Manhattan
    shared_data[warp_id + 8] = warp_mse;        // Second half: MSE
}
__syncthreads();

// Final sum calculation for both metrics
if (tid == 0) {
    float total_manhattan = 0.0f;
    float total_mse = 0.0f;
    int num_warps = blockDim.x / 32;
    for (int i = 0; i < num_warps; i++) {
        total_manhattan += shared_data[i];
        total_mse += shared_data[i + 8];
    }
    manhattan_distances[sample_idx] = total_manhattan;
    mse_losses[sample_idx] = total_mse / feature_dim;
}



7. **FLOAT4 DUAL COMPUTATION**: Integrate both metrics seamlessly with vectorization:
cpp
// Apply dual metric computation to float4 vector
float4 diff;
diff.x = x_val.x - y_val.x;
diff.y = x_val.y - y_val.y;
diff.z = x_val.z - y_val.z;
diff.w = x_val.w - y_val.w;

// Accumulate Manhattan distance
warp_manhattan += fabsf(diff.x) + fabsf(diff.y) + 
                 fabsf(diff.z) + fabsf(diff.w);

// Accumulate MSE loss
warp_mse += diff.x * diff.x + diff.y * diff.y + 
            diff.z * diff.z + diff.w * diff.w;



8. **FAST MODE KERNEL**: Implement atomic operation version:
cpp
// Fast mode kernel with atomic operations
__global__ void manhattan_mse_kernel_fast(
    const float* __restrict__ x,
    const float* __restrict__ y,
    float* __restrict__ manhattan_distances,
    float* __restrict__ mse_losses,
    int batch_size,
    int feature_dim
) {
    // Use shared memory for block-level reduction
    // Use atomicAdd for final accumulation
    // Process both metrics simultaneously
}



9. **EFFICIENT MODE KERNEL**: Implement two-level reduction version:
cpp
// Efficient mode kernel with two-level reduction
__global__ void manhattan_mse_kernel_efficient(
    const float* __restrict__ x,
    const float* __restrict__ y,
    float* __restrict__ manhattan_partial_sums,
    float* __restrict__ mse_partial_sums,
    int batch_size,
    int feature_dim
) {
    // Use shared memory for block-level reduction
    // Write partial sums to global memory
    // CPU handles final reduction
    // Process both metrics simultaneously
}



10. **BLOCK CONFIGURATION**: Use optimal settings for dual metric processing:
    - Block size: 256 threads (8 warps)
    - Shared memory: 16 * sizeof(float) for dual metric results
    - One block per sample for maximum parallelism
    - Elements per warp: (feature_dim + 8 - 1) / 8

11. **PRECISION REQUIREMENTS**: Ensure exact mathematical alignment:
    - Manhattan Distance: Σ|x_i - y_i|
    - MSE Loss: Σ(x_i - y_i)² / feature_dim
    - Use fabsf for absolute value computation
    - Use standard multiplication for squared differences
    - Verify with torch.allclose(rtol=1e-03, atol=1e-6)

12. **FUNCTION SIGNATURE**: The main CUDA function must accept all parameters:
cpp
torch::Tensor manhattan_mse_cuda(
    torch::Tensor x,
    torch::Tensor y,
    std::string mode = "fast"
)



13. **MATHEMATICAL FORMULAS**: Implement exact mathematical operations:
    - Difference: diff = x_i - y_i
    - Manhattan Distance: manhattan_dist = Σ|diff|
    - MSE Loss: mse_loss = Σ(diff²) / feature_dim
    - Both computed from the same diff values

14. **PYTHON CALLING CONVENTION**: The ModelNew forward method must pass parameters correctly:
python
def forward(self, x, y):
    return self.manhattan_mse.manhattan_mse_cuda(x, y, self.mode)



15. **OUTPUT REQUIREMENTS**: Generate both distance metrics:
    - Primary output: MSE losses per sample [batch_size]
    - Secondary output: Manhattan distances per sample [batch_size]
    - Both outputs must match PyTorch reference implementation exactly

16. **PERFORMANCE OPTIMIZATIONS**: Include advanced optimizations:
    - Use fast math optimizations (--use_fast_math)
    - Optimize for compute capability 8.0+ (sm_80)
    - Use -O3 optimization level
    - Avoid bank conflicts in shared memory access
    - Use efficient memory access patterns

17. **ALGORITHM CHOICE**: Prioritize the vectorized fused approach:
    - float4 vectorization is mandatory for this implementation
    - Do NOT implement scalar-only versions
    - The fusion must happen at the CUDA kernel level, not Python level
    - Eliminate all intermediate tensor storage

18. **BOUNDARY HANDLING**: Properly handle non-multiple-of-4 feature dimensions:
    - Use float4 for vectorized processing of main portion
    - Handle remaining elements with scalar processing
    - Ensure no memory access violations
    - Maintain mathematical correctness for both metrics

19. **DUAL METRIC PROPERTIES**: Leverage the relationship between metrics:
    - Both use the same difference values: diff = x - y
    - Manhattan uses absolute values: |diff|
    - MSE uses squared values: diff²
    - Compute both from the same diff to maximize efficiency

20. **MODE SELECTION**: Implement intelligent mode selection:
    - "fast": Use atomic operations for simplicity
    - "efficient": Use two-level reduction for better scalability
    - Default to "fast" mode for backward compatibility
    - Allow runtime mode selection

Here's the target architecture to optimize:

python
import torch
import torch.nn as nn

class Model(nn.Module):
"""
合理优化的PyTorch MSE Loss实现
使用PyTorch内置函数，避免不必要的中间张量创建
"""
def init(self):
super(Model, self).init()

def forward(self, input: torch.Tensor, target: torch.Tensor) -> torch.Tensor:  
    """  
    使用PyTorch内置的mse_loss函数  
    MSE Loss = (input - target)²  

    Args:  
        input (torch.Tensor): 预测值  
        target (torch.Tensor): 真实值  

    Returns:  
        torch.Tensor: MSE Loss标量值  
    """  
    # 直接使用内置函数，让PyTorch处理优化  
    return torch.nn.functional.mse_loss(  
        input,  
        target,  
        reduction='sum'  
    )  
batch_size = 128
num_features = 2000

def get_inputs():
"""
生成合理的测试数据
"""
input_vals = torch.randn(batch_size, num_features)
target_vals = torch.randn(batch_size, num_features)
return [input_vals, target_vals]

def get_init_inputs():
return [] # 没有特殊的初始化输入需求



**EXPECTED OUTPUT STRUCTURE**:
Generate two files:
1. `manhattan_mse_cudacode.py` - Contains ModelNew class with Manhattan+MSE fusion using pure CUDA
2. `manhattan_mse_torchcode.py` - Contains the reference PyTorch implementation with dual metric fusion

**IMPORTANT IMPLEMENTATION REQUIREMENTS**:
1. Use raw pointer access with data_ptr<float>() instead of PackedTensorAccessor
2. Implement TWO distinct CUDA kernels: one with shared memory reduction and atomic operations (fast mode), another with two-level reduction (partial sums + CPU final reduction) for efficient mode
3. Each thread must process multiple elements using stride loop (for (int i = idx; i < size; i += stride))
4. Use shared memory for block-level reduction before atomic operations
5. Use 256 threads per block and limit to max 1024 blocks
6. Include a mode parameter in the main function to select between "fast" (atomic) and "efficient" (two-level) modes
7. Use TORCH_CHECK for input validation
8. Use extern __shared__ float shared_mem[] for shared memory allocation
9. The ModelNew class must accept a mode parameter in __init__ and pass it to the CUDA function
10. Use load_inline with specific compilation flags: -O3, --use_fast_math, -gencode=arch=compute_80,code=sm_80
11. Must implement dual metric fusion (Manhattan + MSE) in a single kernel
12. Must use float4 vectorization for maximum performance
13. Must use warp-level optimization for maximum performance
14. Must handle arbitrary tensor shapes (not just fixed dimensions)
15. Must maintain mathematical precision with PyTorch implementation
16. Expected speedup: 2.2-3.5x over PyTorch baseline
17. Must use fast math optimizations for better performance
18. Must be robust and handle edge cases properly
19. Must use only pure CUDA functions (no PyTorch internal functions)
20. Must use fabsf for absolute value computation
21. Must implement exact mathematical formulas for both Manhattan distance and MSE loss
22. Must generate both distance and loss tensors
23. Must use shared memory efficiently for dual metric warp-level sum reduction
24. Must ensure coalesced memory access patterns
25. Must eliminate intermediate tensor storage for maximum fusion benefits
26. Must implement the complete fusion in a single CUDA kernel
27. Must use float4 vectorization as the primary optimization strategy
28. Must compute both metrics from the same difference values for maximum efficiency
29. Must handle both reduction modes (fast and efficient) appropriately
30. Must return MSE loss as primary output with Manhattan distance as secondary output
